Implement prefetching#9723
Conversation
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
|
See: #3126 (comment) that justifies this PR |
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughSummary by CodeRabbit
WalkthroughAdds ChangesPrefetch helper primitives
Assessment against linked issues
Out-of-scope changes
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 54e03265-c405-4ed7-9a6a-2dda8d2ab5e1
📒 Files selected for processing (6)
cub/benchmarks/bench/reduce/by_key.cucub/cub/agent/agent_reduce_by_key.cuhcub/cub/detail/prefetch.cuhcub/cub/device/dispatch/dispatch_reduce_by_key.cuhcub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuhcub/test/catch2_test_device_reduce_by_key_env.cu
This comment has been minimized.
This comment has been minimized.
gonidelis
force-pushed
the
prefetch_alt_design
branch
from
010f57a to
6c6b182
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: b9f60ed6-bca2-4a64-8026-af6c18a4821f
📒 Files selected for processing (1)
cub/cub/detail/prefetch.cuh
gonidelis
force-pushed
the
prefetch_alt_design
branch
from
233e4e5 to
e299245
Compare
😬 CI Workflow Results🟥 Finished in 1h 18m: Pass: 84%/287 | Total: 1d 20h | Max: 1h 00m | Hits: 96%/204333See results here. |
| static_assert(sizeof(cub::detail::it_value_t<RandomAccessIterator>) == sizeof(T), | ||
| "BlockPrefetch element type T must match the iterator's value type size"); | ||
| const int linear_tid = static_cast<int>(threadIdx.x); | ||
| const unsigned int total_bytes = static_cast<unsigned int>(items_to_prefetch) * unsigned{sizeof(T)}; |
There was a problem hiding this comment.
Critical: we want to prefetch the entire tile not just the data of the first thread:
| const unsigned int total_bytes = static_cast<unsigned int>(items_to_prefetch) * unsigned{sizeof(T)}; | |
| const unsigned int total_bytes = static_cast<unsigned int>(items_to_prefetch) * unsigned{sizeof(T)} * ThreadsPerBlock; |
There was a problem hiding this comment.
i will reject this as we want to express full set of elements to prefetch for edge cases / ragged loads. I will go with the renames you suggested above though.
if i were to apply this suggested fix the Prefetch() call would be problematic:
// full tile: pass items per thread
BlockPrefetch<128, l2>::Prefetch(d_keys_in + tile_offset, 11); // 11 × 128 = 1408 ok
// ragged last tile, 500 items left:
BlockPrefetch<128, l2>::Prefetch(d_keys_in + tile_offset, 4); // 4 * 128 = 512 --> 12 past the end
BlockPrefetch<128, l2>::Prefetch(d_keys_in + tile_offset, 3); // 3 * 128 = 384 --> misses 116
// no integer means "500"…tor params - Use Apache-2.0 WITH LLVM-exception on both new files (BSD-3-Clause is the legacy license; all recently added CUB files use Apache-2.0) - Drop BlockPrefetch's T template parameter and the static_assert that policed it; the element size is derived from the iterator's value type instead - Rename iterator template parameters to It: can_prefetch_from accepts any iterator category and tests for contiguous, so RandomAccessIterator was misleading
- Cover the iterators users actually pass: thrust::device_vector and thrust::universal_vector (rejected by can_prefetch_from, safe no-op) and cuda::buffer's heterogeneous_iterator (contiguous, accepted), per review - Use thrust::no_init for all vectors whose contents get overwritten - Use signed byte/offset math; a non-positive count now trivially no-ops
There was a problem hiding this comment.
Actionable comments posted: 1
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cub/cub/detail/prefetch.cuh (1)
133-135: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick winimportant: keep
l1mapped to L2 until native L1 support is gated.The PR contract says
l1falls back to L2 on current unsupported architectures, but this branch always emitsprefetch.global.L1. Until there is an architecture-gated native path, emit L2 here.if constexpr (PrefetchLevel == LoadPrefetch::l1) { - asm volatile("prefetch.global.L1 [%0];" : : "l"(::cuda::ptx::__as_ptr_gmem(src_ptr + offset)) : "memory"); + asm volatile("prefetch.global.L2 [%0];" : : "l"(::cuda::ptx::__as_ptr_gmem(src_ptr + offset)) : "memory"); }
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 4a3923e5-b623-4845-80b7-4ab8e42f07db
📒 Files selected for processing (2)
cub/cub/detail/prefetch.cuhcub/test/catch2_test_block_prefetch.cu
| //! @tparam ThreadsPerBlock Number of threads in the block cooperating on the hints. | ||
| //! @tparam PrefetchLevel Which cache level to target (see ``LoadPrefetch``); ``none`` makes ``Prefetch()`` a no-op. | ||
| //! @tparam PrefetchStride Byte stride between successive hints; one hint per cache line (default 128 B). | ||
| template <int ThreadsPerBlock, LoadPrefetch PrefetchLevel = LoadPrefetch::l2, int PrefetchStride = 128> |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
important: reject non-positive block/stride configurations.
PrefetchStride <= 0 or ThreadsPerBlock <= 0 makes the loop on Line 130 stop progressing or address backwards. Add compile-time guards on the template.
template <int ThreadsPerBlock, LoadPrefetch PrefetchLevel = LoadPrefetch::l2, int PrefetchStride = 128>
struct BlockPrefetch
{
+ static_assert(ThreadsPerBlock > 0, "ThreadsPerBlock must be positive");
+ static_assert(PrefetchStride > 0, "PrefetchStride must be positive");
+As per path instructions, focus on algorithm correctness, memory access safety, and CUDA synchronization/performance risks.
Also applies to: 130-130
Source: Path instructions
2 participants
fixes #9588 with a new design.
Adds
cub::detail::BlockPrefetch, a standalone block-scope facility for cooperatively emitting global-memory tile-prefetch hints.Design-only PR: it lands the facility and nothing else, no algorithm is wired to it yet (net diff is a single new header).
What this adds
cub::detail::LoadPrefetch- enum of prefetch targets:none,l2(prefetch.global.L2), l1 (prefetch.global.L1, falls back to L2 where real L1 prefetch is unavailable), andbulk_l2(TMAcp.async.bulk.prefetchinto L2; requires SM90+, no-op on older arch).cub::detail::BlockPrefetch<T, ThreadsPerBlock, Prefetch, PrefetchStride>- a collective with a single staticPrefetch(it, items_to_prefetch). The block's threads cooperatively stride the tile, one hint per cache line (or one TMA bulk prefetch forbulk_l2). A no-op for none and for iterators that can't yield a raw address (e.g.CacheModifiedInputIterator), so call sites need no enablement check.Everything is under detail, no public-API or stability commitment.